/**
 * Copyright (c) 2024 Huawei Technologies Co., Ltd.
 * This file is a part of the CANN Open Software.
 * Licensed under CANN Open Software License Agreement Version 1.0 (the "License").
 * Please refer to the License for details. You may not use this file except in compliance with the License.
 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
 * See LICENSE in the root of the software repository for the full text of the License.
 */

#include "../kernel_impl/softmaxgrad_kernel.h"
using namespace AscendC;

__aicore__ inline void CopyTiling(VecTiling* tiling, GM_ADDR tilingGM)
{
    uint32_t* ptr = reinterpret_cast<uint32_t*>(tiling);
    auto tiling32 = reinterpret_cast<__gm__ uint32_t*>(tilingGM);

    for (int i = 0; i < sizeof(VecTiling) / sizeof(uint32_t); i++, ptr++) {
        *ptr = *(tiling32 + i);
    }
    return;
}

extern "C" __global__ __aicore__ void SoftmaxgradCustom(GM_ADDR x, GM_ADDR y, GM_ADDR res_out, GM_ADDR workspace,
    GM_ADDR tiling)
{
    if ASCEND_IS_AIC {
        return;
    }
    VecTiling tilingData;
    CopyTiling(&tilingData, tiling);
    KernelSoftmax op;
    op.Init(x, y, res_out, tilingData);
    op.Process();
}

#ifndef ASCENDC_CPU_DEBUG
// call of kernel function
void SoftmaxgradCustomDo(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* res_out,
    uint8_t* workspace, uint8_t* tiling)
{
    SoftmaxgradCustom<<<blockDim, l2ctrl, stream>>>(x, y, res_out, workspace, tiling);
}
#endif